<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>MC-GPU: MC-GPU_kernel_v1.3.cu File Reference</title>

<link href="tabs.css" rel="stylesheet" type="text/css"/>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>


</head>
<body>
<div id="top"><!-- do not remove this div! -->


<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  
  
  <td style="padding-left: 0.5em;">
   <div id="projectname">MC-GPU
   
   </div>
   
  </td>
  
  
  
 </tr>
 </tbody>
</table>
</div>

<!-- Generated by Doxygen 1.7.6.1 -->
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="annotated.html"><span>Classes</span></a></li>
      <li class="current"><a href="files.html"><span>Files</span></a></li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="files.html"><span>File&#160;List</span></a></li>
      <li><a href="globals.html"><span>File&#160;Members</span></a></li>
    </ul>
  </div>
</div>
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
  initNavTree('MC-GPU__kernel__v1_83_8cu.html','');
</script>
<div id="doc-content">
<div class="header">
  <div class="summary">
<a href="#define-members">Defines</a> &#124;
<a href="#func-members">Functions</a>  </div>
  <div class="headertitle">
<div class="title">MC-GPU_kernel_v1.3.cu File Reference</div>  </div>
</div><!--header-->
<div class="contents">

<p>Definition of the CUDA GPU kernel for the simulation of x ray tracks in a voxelized geometry.  
<a href="#details">More...</a></p>

<p><a href="MC-GPU__kernel__v1_83_8cu_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr><td colspan="2"><h2><a name="define-members"></a>
Defines</h2></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a3ab75c5bf72afd4a9b0255f61bd95704">LEAP_DISTANCE</a>&#160;&#160;&#160;256</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Upper limit of the number of random values sampled in a single track.  <a href="#a3ab75c5bf72afd4a9b0255f61bd95704"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#afd7d5320ad79ad99c8d230f3a730141d">a1_RANECU</a>&#160;&#160;&#160;40014</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Multipliers and moduli for the two MLCG in RANECU.  <a href="#afd7d5320ad79ad99c8d230f3a730141d"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#aa7d2cfcd3830e91736ca79f4bfdfb5d2">m1_RANECU</a>&#160;&#160;&#160;2147483563</td></tr>
<tr><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a527a42aca15fac6a1cab1116a72eb8ca">a2_RANECU</a>&#160;&#160;&#160;40692</td></tr>
<tr><td class="memItemLeft" align="right" valign="top">#define&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a9872bd70bdcc27f56bc025f291b7d6c2">m2_RANECU</a>&#160;&#160;&#160;2147483399</td></tr>
<tr><td colspan="2"><h2><a name="func-members"></a>
Functions</h2></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#ab1e2e311c809fe74d093498e53c9c954">track_particles</a> (int history_batch, int histories_per_thread, int num_p, int seed_input, unsigned long long int *image, <a class="el" href="structulonglong2.html">ulonglong2</a> *voxels_Edep, <a class="el" href="structfloat2.html">float2</a> *voxel_mat_dens, <a class="el" href="structfloat2.html">float2</a> *mfp_Woodcock_table, <a class="el" href="structfloat3.html">float3</a> *mfp_table_a, <a class="el" href="structfloat3.html">float3</a> *mfp_table_b, struct <a class="el" href="structrayleigh__struct.html">rayleigh_struct</a> *rayleigh_table, struct <a class="el" href="structcompton__struct.html">compton_struct</a> *compton_table, struct <a class="el" href="structdetector__struct.html">detector_struct</a> *detector_data_array, struct <a class="el" href="structsource__struct.html">source_struct</a> *source_data_array, <a class="el" href="structulonglong2.html">ulonglong2</a> *materials_dose)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Initialize the image array, ie, set all pixels to zero Essentially, this function has the same effect as the command: "cutilSafeCall(cudaMemcpy(image_device, image, image_bytes, cudaMemcpyHostToDevice))";.  <a href="#ab1e2e311c809fe74d093498e53c9c954"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#aac6cb77943874a858ac5a79cbd58f1e7">tally_voxel_energy_deposition</a> (float *Edep, <a class="el" href="structshort3.html">short3</a> *voxel_coord, <a class="el" href="structulonglong2.html">ulonglong2</a> *voxels_Edep)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Tally the dose deposited in the voxels.  <a href="#aac6cb77943874a858ac5a79cbd58f1e7"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#aac14e3b20e073c6b169f2bd5c23da281">tally_image</a> (float *energy, <a class="el" href="structfloat3.html">float3</a> *position, <a class="el" href="structfloat3.html">float3</a> *direction, signed char *scatter_state, unsigned long long int *image, struct <a class="el" href="structsource__struct.html">source_struct</a> *source_data_SHARED, struct <a class="el" href="structdetector__struct.html">detector_struct</a> *detector_data_SHARED)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Tally a radiographic projection image.  <a href="#aac14e3b20e073c6b169f2bd5c23da281"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a6807df36f6f2501a9f57ea4fce81b4f2">source</a> (<a class="el" href="structfloat3.html">float3</a> *position, <a class="el" href="structfloat3.html">float3</a> *direction, float *energy, <a class="el" href="structint2.html">int2</a> *seed, int *absvox, struct <a class="el" href="structsource__struct.html">source_struct</a> *source_data_SHARED, struct <a class="el" href="structdetector__struct.html">detector_struct</a> *detector_data_SHARED)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Source that creates primary x rays, according to the defined source model.  <a href="#a6807df36f6f2501a9f57ea4fce81b4f2"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a047931bac84f3bce50dc13061371e765">move_to_bbox</a> (<a class="el" href="structfloat3.html">float3</a> *position, <a class="el" href="structfloat3.html">float3</a> *direction, <a class="el" href="structfloat3.html">float3</a> size_bbox, int *intersection_flag)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Functions that moves a particle inside the voxelized geometry bounding box.  <a href="#a047931bac84f3bce50dc13061371e765"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#ae7bd0b93b4d5eedb15822fa0ef622b48">init_PRNG</a> (int history_batch, int histories_per_thread, int seed_input, <a class="el" href="structint2.html">int2</a> *seed)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Initialize the pseudo-random number generator (PRNG) RANECU to a position far away from the previous history (leap frog technique).  <a href="#ae7bd0b93b4d5eedb15822fa0ef622b48"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">int&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#ae9061f4314474bf4415200a5cb115397">abMODm</a> (int m, int a, int s)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Calculate "(a1*a2) MOD m" with 32-bit integers and avoiding the possible overflow, using the Russian Peasant approach modulo m and the approximate factoring method, as described in: L'Ecuyer and Cote, ACM Trans.  <a href="#ae9061f4314474bf4415200a5cb115397"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">float&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a3b57f2faa5765ceba5d352641ca72a9f">ranecu</a> (<a class="el" href="structint2.html">int2</a> *seed)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Pseudo-random number generator (PRNG) RANECU returning a float value (single precision version).  <a href="#a3b57f2faa5765ceba5d352641ca72a9f"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">double&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a74c4894502584f8fb113e04b759a6c55">ranecu_double</a> (<a class="el" href="structint2.html">int2</a> *seed)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Pseudo-random number generator (PRNG) RANECU returning a double value.  <a href="#a74c4894502584f8fb113e04b759a6c55"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">int&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a562dba755f1abdb20b1088d3121f4326">locate_voxel</a> (<a class="el" href="structfloat3.html">float3</a> *position, <a class="el" href="structshort3.html">short3</a> *voxel_coord)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Find the voxel that contains the current position.  <a href="#a562dba755f1abdb20b1088d3121f4326"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a94ad9ed794eb8fd696e7e357ae391765">rotate_double</a> (<a class="el" href="structfloat3.html">float3</a> *direction, double costh, double phi)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Rotates a vector; the rotation is specified by giving the polar and azimuthal angles in the "self-frame", as determined by the vector to be rotated.  <a href="#a94ad9ed794eb8fd696e7e357ae391765"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a5cafd4ebb3d155746b16ca7d231ecad1">GRAa</a> (float *energy, double *costh_Rayleigh, int *mat, float *pmax_current, <a class="el" href="structint2.html">int2</a> *seed, struct <a class="el" href="structrayleigh__struct.html">rayleigh_struct</a> *cgra)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Sample a Rayleigh interaction using the sampling algorithm used in PENELOPE 2006.  <a href="#a5cafd4ebb3d155746b16ca7d231ecad1"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a6f212862e7fba1553ec9495f05443b2c">GCOa</a> (float *energy, double *costh_Compton, int *mat, <a class="el" href="structint2.html">int2</a> *seed, struct <a class="el" href="structcompton__struct.html">compton_struct</a> *cgco_SHARED)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Random sampling of incoherent (Compton) scattering of photons, using the sampling algorithm from PENELOPE 2006: Relativistic impulse approximation with analytical one-electron Compton profiles.  <a href="#a6f212862e7fba1553ec9495f05443b2c"></a><br/></td></tr>
<tr><td class="memItemLeft" align="right" valign="top">void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a45a0719009ba55cca7af72f08f1090fb">tally_materials_dose</a> (float *Edep, int *material, <a class="el" href="structulonglong2.html">ulonglong2</a> *materials_dose)</td></tr>
<tr><td class="mdescLeft">&#160;</td><td class="mdescRight">Tally the depose deposited inside each material.  <a href="#a45a0719009ba55cca7af72f08f1090fb"></a><br/></td></tr>
</table>
<hr/><a name="details" id="details"></a><h2>Detailed Description</h2>
<div class="textblock"><p>Definition of the CUDA GPU kernel for the simulation of x ray tracks in a voxelized geometry. </p>
<p>This kernel has been optimized to yield a good performance in the GPU but can still be compiled in the CPU without problems. All the CUDA especific commands are enclosed in pre-processor directives that are skipped if the parameter "USING_CUDA" is not defined at compilation time. </p>
<dl class="author"><dt><b>Author:</b></dt><dd>Andreu Badal (<a href="mailto:Andreu.Badal-Soler@fda.hhs.gov">Andreu.Badal-Soler@fda.hhs.gov</a>) </dd></dl>
<dl class="date"><dt><b>Date:</b></dt><dd>2012/12/12 </dd></dl>

<p>Definition in file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>
</div><hr/><h2>Define Documentation</h2>
<a class="anchor" id="afd7d5320ad79ad99c8d230f3a730141d"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::a1_RANECU" ref="afd7d5320ad79ad99c8d230f3a730141d" args="" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">#define <a class="el" href="MC-GPU__kernel__v1_83_8cu.html#afd7d5320ad79ad99c8d230f3a730141d">a1_RANECU</a>&#160;&#160;&#160;40014</td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Multipliers and moduli for the two MLCG in RANECU. </p>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00814">814</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>, and <a class="el" href="MC-GPU__v1_83_8cu_source.html#l03356">update_seed_PRNG()</a>.</p>

</div>
</div>
<a class="anchor" id="a527a42aca15fac6a1cab1116a72eb8ca"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::a2_RANECU" ref="a527a42aca15fac6a1cab1116a72eb8ca" args="" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">#define <a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a527a42aca15fac6a1cab1116a72eb8ca">a2_RANECU</a>&#160;&#160;&#160;40692</td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00816">816</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>.</p>

</div>
</div>
<a class="anchor" id="a3ab75c5bf72afd4a9b0255f61bd95704"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::LEAP_DISTANCE" ref="a3ab75c5bf72afd4a9b0255f61bd95704" args="" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">#define <a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a3ab75c5bf72afd4a9b0255f61bd95704">LEAP_DISTANCE</a>&#160;&#160;&#160;256</td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Upper limit of the number of random values sampled in a single track. </p>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00812">812</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>, and <a class="el" href="MC-GPU__v1_83_8cu_source.html#l03356">update_seed_PRNG()</a>.</p>

</div>
</div>
<a class="anchor" id="aa7d2cfcd3830e91736ca79f4bfdfb5d2"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::m1_RANECU" ref="aa7d2cfcd3830e91736ca79f4bfdfb5d2" args="" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">#define <a class="el" href="MC-GPU__kernel__v1_83_8cu.html#aa7d2cfcd3830e91736ca79f4bfdfb5d2">m1_RANECU</a>&#160;&#160;&#160;2147483563</td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00815">815</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>, and <a class="el" href="MC-GPU__v1_83_8cu_source.html#l03356">update_seed_PRNG()</a>.</p>

</div>
</div>
<a class="anchor" id="a9872bd70bdcc27f56bc025f291b7d6c2"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::m2_RANECU" ref="a9872bd70bdcc27f56bc025f291b7d6c2" args="" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">#define <a class="el" href="MC-GPU__kernel__v1_83_8cu.html#a9872bd70bdcc27f56bc025f291b7d6c2">m2_RANECU</a>&#160;&#160;&#160;2147483399</td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00817">817</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>.</p>

</div>
</div>
<hr/><h2>Function Documentation</h2>
<a class="anchor" id="ae9061f4314474bf4415200a5cb115397"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::abMODm" ref="ae9061f4314474bf4415200a5cb115397" args="(int m, int a, int s)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">int <a class="el" href="MC-GPU__v1_83_8h.html#ae9061f4314474bf4415200a5cb115397">abMODm</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>m</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>a</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>s</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Calculate "(a1*a2) MOD m" with 32-bit integers and avoiding the possible overflow, using the Russian Peasant approach modulo m and the approximate factoring method, as described in: L'Ecuyer and Cote, ACM Trans. </p>
<p>Math. Soft. 17 (1991).</p>
<p>This function has been adapted from "seedsMLCG.f", see: Badal and Sempau, Computer Physics Communications 175 (2006)</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">m,a,s</td><td>MLCG parameters </td></tr>
  </table>
  </dd>
</dl>
<dl class="return"><dt><b>Returns:</b></dt><dd>(a1*a2) MOD m </dd></dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00919">919</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>, and <a class="el" href="MC-GPU__v1_83_8cu_source.html#l03356">update_seed_PRNG()</a>.</p>

</div>
</div>
<a class="anchor" id="a6f212862e7fba1553ec9495f05443b2c"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::GCOa" ref="a6f212862e7fba1553ec9495f05443b2c" args="(float *energy, double *costh_Compton, int *mat, int2 *seed, struct compton_struct *cgco_SHARED)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#a6f212862e7fba1553ec9495f05443b2c">GCOa</a> </td>
          <td>(</td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>energy</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">double *&#160;</td>
          <td class="paramname"><em>costh_Compton</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int *&#160;</td>
          <td class="paramname"><em>mat</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structcompton__struct.html">compton_struct</a> *&#160;</td>
          <td class="paramname"><em>cgco_SHARED</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Random sampling of incoherent (Compton) scattering of photons, using the sampling algorithm from PENELOPE 2006: Relativistic impulse approximation with analytical one-electron Compton profiles. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">energy</td><td>incident and final photon energy (eV) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">costh_Compton</td><td>cosine of the polar scattering angle </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">material</td><td>Current voxel material </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">seed</td><td>RANECU PRNG seed </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01287">1287</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00241">compton_struct::fco</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00241">compton_struct::fj0</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00062">MAX_MATERIALS</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00063">MAX_SHELLS</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00093">max_value</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00094">min_value</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00244">compton_struct::noscco</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00965">ranecu()</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00241">compton_struct::uico</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a5cafd4ebb3d155746b16ca7d231ecad1"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::GRAa" ref="a5cafd4ebb3d155746b16ca7d231ecad1" args="(float *energy, double *costh_Rayleigh, int *mat, float *pmax_current, int2 *seed, struct rayleigh_struct *cgra)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#a5cafd4ebb3d155746b16ca7d231ecad1">GRAa</a> </td>
          <td>(</td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>energy</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">double *&#160;</td>
          <td class="paramname"><em>costh_Rayleigh</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int *&#160;</td>
          <td class="paramname"><em>mat</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>pmax_current</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structrayleigh__struct.html">rayleigh_struct</a> *&#160;</td>
          <td class="paramname"><em>cgra</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Sample a Rayleigh interaction using the sampling algorithm used in PENELOPE 2006. </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">energy</td><td>Particle energy (not modified with Rayleigh) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">costh_Rayleigh</td><td>Cosine of the angular deflection </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">material</td><td>Current voxel material </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01181">1181</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00254">rayleigh_struct::aco</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00254">rayleigh_struct::bco</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00259">rayleigh_struct::itlco</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00259">rayleigh_struct::ituco</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00094">min_value</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00064">NP_RAYLEIGH</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00254">rayleigh_struct::pco</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00995">ranecu_double()</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00254">rayleigh_struct::xco</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="ae7bd0b93b4d5eedb15822fa0ef622b48"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::init_PRNG" ref="ae7bd0b93b4d5eedb15822fa0ef622b48" args="(int history_batch, int histories_per_thread, int seed_input, int2 *seed)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#ae7bd0b93b4d5eedb15822fa0ef622b48">init_PRNG</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>history_batch</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>histories_per_thread</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>seed_input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Initialize the pseudo-random number generator (PRNG) RANECU to a position far away from the previous history (leap frog technique). </p>
<p>Each calculated seed initiates a consecutive and disjoint sequence of pseudo-random numbers with length LEAP_DISTANCE, that can be used to in a parallel simulation (Sequence Splitting parallelization method). The basic equation behind the algorithm is: S(i+j) = (a**j * S(i)) MOD m = [(a**j MOD m)*S(i)] MOD m , which is described in: P L'Ecuyer, Commun. ACM 31 (1988) p.742</p>
<p>This function has been adapted from "seedsMLCG.f", see: A Badal and J Sempau, Computer Physics Communications 175 (2006) p. 440-450</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">history</td><td>Particle bach number. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">seed_input</td><td>Initial PRNG seed input (used to initiate both MLCGs in RANECU). </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">seed</td><td>Initial PRNG seeds for the present history. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">841</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00814">a1_RANECU</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00816">a2_RANECU</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00919">abMODm()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00812">LEAP_DISTANCE</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00815">m1_RANECU</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00817">m2_RANECU</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::x</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::y</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a562dba755f1abdb20b1088d3121f4326"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::locate_voxel" ref="a562dba755f1abdb20b1088d3121f4326" args="(float3 *position, short3 *voxel_coord)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">int <a class="el" href="MC-GPU__v1_83_8h.html#a562dba755f1abdb20b1088d3121f4326">locate_voxel</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>position</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structshort3.html">short3</a> *&#160;</td>
          <td class="paramname"><em>voxel_coord</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Find the voxel that contains the current position. </p>
<p>Report the voxel absolute index and the x,y,z indices. The structure containing the voxel number and size is read from CONSTANT memory.</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">position</td><td>Particle position </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">voxel_coord</td><td>Pointer to three integer values (short3*) that will store the x,y and z voxel indices. </td></tr>
  </table>
  </dd>
</dl>
<dl class="return"><dt><b>Returns:</b></dt><dd>Returns "absvox", the voxel number where the particle is located (negative if position outside the voxel bbox). </dd></dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01033">1033</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00085">EPS_SOURCE</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00216">voxel_struct::inv_voxel_size</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00215">voxel_struct::num_voxels</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00216">voxel_struct::size_bbox</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00317">voxel_data_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00131">int3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00131">int3::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a047931bac84f3bce50dc13061371e765"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::move_to_bbox" ref="a047931bac84f3bce50dc13061371e765" args="(float3 *position, float3 *direction, float3 size_bbox, int *intersection_flag)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#a047931bac84f3bce50dc13061371e765">move_to_bbox</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>position</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>direction</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a>&#160;</td>
          <td class="paramname"><em>size_bbox</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int *&#160;</td>
          <td class="paramname"><em>intersection_flag</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Functions that moves a particle inside the voxelized geometry bounding box. </p>
<p>An EPSILON distance is added to make sure the particles will be clearly inside the bbox, not exactly on the surface.</p>
<p>This algorithm makes the following assumtions:</p>
<ul>
<li>The back lower vertex of the voxel bounding box is always located at the origin: (x0,y0,z0)=(0,0,0).</li>
<li>The initial value of "position" corresponds to the focal spot location.</li>
<li>When a ray is not pointing towards the bbox plane that it should cross according to the sign of the direction, I assign a distance to the intersection =0 instead of the real negative distance. The wall that will be crossed to enter the bbox is always the furthest and therefore a 0 distance will never be used except in the case of a ray starting inside the bbox or outside the bbox and not pointing to any of the 3 planes. In this situation the ray will be transported a 0 distance, meaning that it will stay at the focal spot.</li>
</ul>
<p>(Interesting information on ray-box intersection: <a href="http://tog.acm.org/resources/GraphicsGems/gems/RayBox.c">http://tog.acm.org/resources/GraphicsGems/gems/RayBox.c</a>)</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">position</td><td>Particle position: initially set to the focal spot, returned transported inside the voxel bbox. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">direction</td><td>Sampled particle direction (cosine vectors). </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">intersection_flag</td><td>Set to &lt;0 if particle outside bbox and will not cross the voxels, not changed otherwise. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">size_bbox</td><td>Size of the bounding box. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00714">714</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00085">EPS_SOURCE</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00087">NEG_EPS_SOURCE</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00090">NEG_INF</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00626">source()</a>.</p>

</div>
</div>
<a class="anchor" id="a3b57f2faa5765ceba5d352641ca72a9f"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::ranecu" ref="a3b57f2faa5765ceba5d352641ca72a9f" args="(int2 *seed)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">float <a class="el" href="MC-GPU__v1_83_8h.html#a3b57f2faa5765ceba5d352641ca72a9f">ranecu</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em></td><td>)</td>
          <td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Pseudo-random number generator (PRNG) RANECU returning a float value (single precision version). </p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">seed</td><td>PRNG seed (seed kept in the calling function and updated here). </td></tr>
  </table>
  </dd>
</dl>
<dl class="return"><dt><b>Returns:</b></dt><dd>PRN double value in the open interval (0,1) </dd></dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00965">965</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::x</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::y</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01287">GCOa()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00626">source()</a>, and <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a74c4894502584f8fb113e04b759a6c55"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::ranecu_double" ref="a74c4894502584f8fb113e04b759a6c55" args="(int2 *seed)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">double <a class="el" href="MC-GPU__v1_83_8h.html#a74c4894502584f8fb113e04b759a6c55">ranecu_double</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em></td><td>)</td>
          <td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Pseudo-random number generator (PRNG) RANECU returning a double value. </p>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00995">995</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::x</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::y</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01181">GRAa()</a>, and <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a94ad9ed794eb8fd696e7e357ae391765"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::rotate_double" ref="a94ad9ed794eb8fd696e7e357ae391765" args="(float3 *direction, double costh, double phi)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#acc5186c685730d287039c737445bbe43">rotate_double</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>direction</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">double&#160;</td>
          <td class="paramname"><em>costh</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">double&#160;</td>
          <td class="paramname"><em>phi</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Rotates a vector; the rotation is specified by giving the polar and azimuthal angles in the "self-frame", as determined by the vector to be rotated. </p>
<p>This function is a literal translation from Fortran to C of PENELOPE (v. 2006) subroutine "DIRECT".</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">(u,v,w)</td><td>input vector (=d) in the lab. frame; returns the rotated vector components in the lab. frame </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">costh</td><td>cos(theta), angle between d before and after turn </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">phi</td><td>azimuthal angle (rad) turned by d in its self-frame </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01103">1103</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a6807df36f6f2501a9f57ea4fce81b4f2"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::source" ref="a6807df36f6f2501a9f57ea4fce81b4f2" args="(float3 *position, float3 *direction, float *energy, int2 *seed, int *absvox, struct source_struct *source_data_SHARED, struct detector_struct *detector_data_SHARED)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#a6807df36f6f2501a9f57ea4fce81b4f2">source</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>position</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>direction</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>energy</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structint2.html">int2</a> *&#160;</td>
          <td class="paramname"><em>seed</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int *&#160;</td>
          <td class="paramname"><em>absvox</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structsource__struct.html">source_struct</a> *&#160;</td>
          <td class="paramname"><em>source_data_SHARED</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structdetector__struct.html">detector_struct</a> *&#160;</td>
          <td class="paramname"><em>detector_data_SHARED</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Source that creates primary x rays, according to the defined source model. </p>
<p>The particles are automatically moved to the surface of the voxel bounding box, to start the tracking inside a real material. If the sampled particle do not enter the voxels, it is init in the focal spot and the main program will check if it arrives at the detector or not.</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">source_data</td><td>Structure describing the source. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">source_energy_data_CONST</td><td>Global variable in constant memory space describing the source energy spectrum. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">position</td><td>Initial particle position (particle transported inside the voxel bbox). </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">direction</td><td>Sampled particle direction (cosine vectors). </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">energy</td><td>Sampled energy of the new x ray. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">seed</td><td>Current seed of the random number generator, requiered to sample the movement direction. </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">absvox</td><td>Set to &lt;0 if primary particle will not cross the voxels, not changed otherwise (&gt;0). </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00626">626</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::cos_theta_low</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::D_cos_theta</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::D_phi</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00178">source_energy_struct::espc</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00180">source_energy_struct::espc_alias</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00178">source_energy_struct::espc_cutoff</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::max_height_at_y1cm</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00714">move_to_bbox()</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00177">source_energy_struct::num_bins_espc</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::phi_low</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00159">source_struct::position</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00965">ranecu()</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00161">source_struct::rot_fan</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00203">detector_struct::rotation_flag</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00216">voxel_struct::size_bbox</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00341">source_energy_data_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00317">voxel_data_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="aac14e3b20e073c6b169f2bd5c23da281"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::tally_image" ref="aac14e3b20e073c6b169f2bd5c23da281" args="(float *energy, float3 *position, float3 *direction, signed char *scatter_state, unsigned long long int *image, struct source_struct *source_data_SHARED, struct detector_struct *detector_data_SHARED)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#aac14e3b20e073c6b169f2bd5c23da281">tally_image</a> </td>
          <td>(</td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>energy</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>position</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>direction</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">signed char *&#160;</td>
          <td class="paramname"><em>scatter_state</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">unsigned long long int *&#160;</td>
          <td class="paramname"><em>image</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structsource__struct.html">source_struct</a> *&#160;</td>
          <td class="paramname"><em>source_data_SHARED</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structdetector__struct.html">detector_struct</a> *&#160;</td>
          <td class="paramname"><em>detector_data_SHARED</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Tally a radiographic projection image. </p>
<p>This function is called whenever a particle escapes the voxelized volume. The code checks if the particle would arrive at the detector if it kept moving in a straight line after exiting the voxels (assuming vacuum enclosure). An ideal image formation model is implemented: each pixel counts the total energy of the x rays that enter the pixel (100% detection efficiency for any energy). The image due to primaries and different kinds of scatter is tallied separately.</p>
<p>In the GPU, and atomicAdd() function is used to make sure that multiple threads do not update the same pixel at the same time, which would result in a lose of information. Since the atomicAdd function is only available for 'unsigned long long int' data, the float pixel values are scaled by a factor "SCALE_eV" defined in the header file (eg, #define SCALE_eV 10000.0f) and stored as unsigned long long integers in main memory.</p>
<p>WARNING! If the total tallied signal (for all particles) is larger than "1.8e19/SCALE_eV", there will be a bit overflow and the value will be reset to 0 giving bogus results.</p>
<p>WARNING! The detector plane should be located outside the voxels bounding box. However, since the particles are moved outside the bbox in the last step, they could cross the detector plane anyway. If the particles are less than 2.0 cm behind the detector, they are moved back and detected. Therefore the detector can be a few cm inside the bbox and still work. If the Woodcock mean free path is larger than the distance from the bbox to the detector, we may lose some particles behind the detector!</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">energy</td><td>X-ray energy </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">position</td><td>Particle position </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">direction</td><td>Particle direction (cosine vectors) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">scatter_state</td><td>Flag marking primaries, single Compton, single Rayleigh or multiple scattered radiation </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">image</td><td>Integer array containing the image, ie, the pixel values (in tenths of meV) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00482">482</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00195">detector_struct::center</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00195">detector_struct::corner_min_rotated_to_Y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00159">source_struct::direction</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00197">detector_struct::inv_pixel_size_X</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00197">detector_struct::inv_pixel_size_Z</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00202">detector_struct::num_pixels</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00197">detector_struct::rot_inv</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00203">detector_struct::rotation_flag</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00079">SCALE_eV</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00203">detector_struct::total_num_pixels</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00130">int2::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="a45a0719009ba55cca7af72f08f1090fb"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::tally_materials_dose" ref="a45a0719009ba55cca7af72f08f1090fb" args="(float *Edep, int *material, ulonglong2 *materials_dose)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#a45a0719009ba55cca7af72f08f1090fb">tally_materials_dose</a> </td>
          <td>(</td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>Edep</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int *&#160;</td>
          <td class="paramname"><em>material</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structulonglong2.html">ulonglong2</a> *&#160;</td>
          <td class="paramname"><em>materials_dose</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Tally the depose deposited inside each material. </p>
<p>This function is called whenever a particle suffers a Compton or photoelectric interaction. The energy released in each interaction is added and later in the report function the total deposited energy is divided by the total mass of the material in the voxelized object to get the dose. This naturally accounts for multiple densities for voxels with the same material (not all voxels have same mass). Electrons are not transported in MC-GPU and therefore we are approximating that the dose is equal to the KERMA (energy released by the photons alone). This approximation is acceptable when there is electronic equilibrium and when the range of the secondary electrons is shorter than the organ size.</p>
<p>The function uses atomic functions for a thread-safe access to the GPU memory. We can check if this tally was disabled in the input file checking if the array materials_dose was allocated in the GPU (disabled if pointer = NULL).</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">Edep</td><td>Energy deposited in the interaction </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">material</td><td>Current material id number </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">materials_dose</td><td><a class="el" href="structulonglong2.html">ulonglong2</a> array storing the mateials dose [in eV/g] and dose^2 (ie, uncertainty). </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01547">1547</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00079">SCALE_eV</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00137">ulonglong2::x</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00137">ulonglong2::y</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="aac6cb77943874a858ac5a79cbd58f1e7"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::tally_voxel_energy_deposition" ref="aac6cb77943874a858ac5a79cbd58f1e7" args="(float *Edep, short3 *voxel_coord, ulonglong2 *voxels_Edep)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#abf1491ec13d6068d2753516a3721cfe2">tally_voxel_energy_deposition</a> </td>
          <td>(</td>
          <td class="paramtype">float *&#160;</td>
          <td class="paramname"><em>Edep</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structshort3.html">short3</a> *&#160;</td>
          <td class="paramname"><em>voxel_coord</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structulonglong2.html">ulonglong2</a> *&#160;</td>
          <td class="paramname"><em>voxels_Edep</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td><code> [inline]</code></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Tally the dose deposited in the voxels. </p>
<p>This function is called whenever a particle suffers a Compton or photoelectric interaction. It is not necessary to call this function if the dose tally was disabled in the input file (ie, dose_ROI_x_max_CONST &lt; 0). Electrons are not transported in MC-GPU and therefore we are approximating that the dose is equal to the KERMA (energy released by the photons alone). This approximation is acceptable when there is electronic equilibrium and when the range of the secondary electrons is shorter than the voxel size. Usually the doses will be acceptable for photon energies below 1 MeV. The dose estimates may not be accurate at the interface of low density volumes.</p>
<p>We need to use atomicAdd() in the GPU to prevent that multiple threads update the same voxel at the same time, which would result in a lose of information. This is very improbable when using a large number of voxels but gives troubles with a simple geometries with few voxels (in this case the atomicAdd will slow down the code because threads will update the voxel dose secuentially).</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">Edep</td><td>Energy deposited in the interaction </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">voxel_coord</td><td>Voxel coordinates, needed to check if particle located inside the input region of interest (ROI) </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">voxels_Edep</td><td><a class="el" href="structulonglong2.html">ulonglong2</a> array containing the 3D voxel dose and dose^2 (ie, uncertainty) as unsigned integers scaled by SCALE_eV. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00418">418</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_x_max_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_x_min_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_y_max_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_y_min_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_z_max_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_z_min_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00079">SCALE_eV</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00137">ulonglong2::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00137">ulonglong2::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00136">short3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">track_particles()</a>.</p>

</div>
</div>
<a class="anchor" id="ab1e2e311c809fe74d093498e53c9c954"></a><!-- doxytag: member="MC&#45;GPU_kernel_v1.3.cu::track_particles" ref="ab1e2e311c809fe74d093498e53c9c954" args="(int history_batch, int histories_per_thread, int num_p, int seed_input, unsigned long long int *image, ulonglong2 *voxels_Edep, float2 *voxel_mat_dens, float2 *mfp_Woodcock_table, float3 *mfp_table_a, float3 *mfp_table_b, struct rayleigh_struct *rayleigh_table, struct compton_struct *compton_table, struct detector_struct *detector_data_array, struct source_struct *source_data_array, ulonglong2 *materials_dose)" -->
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">void <a class="el" href="MC-GPU__v1_83_8h.html#ab1e2e311c809fe74d093498e53c9c954">track_particles</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>history_batch</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>histories_per_thread</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>num_p</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>seed_input</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">unsigned long long int *&#160;</td>
          <td class="paramname"><em>image</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structulonglong2.html">ulonglong2</a> *&#160;</td>
          <td class="paramname"><em>voxels_Edep</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat2.html">float2</a> *&#160;</td>
          <td class="paramname"><em>voxel_mat_dens</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat2.html">float2</a> *&#160;</td>
          <td class="paramname"><em>mfp_Woodcock_table</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>mfp_table_a</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structfloat3.html">float3</a> *&#160;</td>
          <td class="paramname"><em>mfp_table_b</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structrayleigh__struct.html">rayleigh_struct</a> *&#160;</td>
          <td class="paramname"><em>rayleigh_table</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structcompton__struct.html">compton_struct</a> *&#160;</td>
          <td class="paramname"><em>compton_table</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structdetector__struct.html">detector_struct</a> *&#160;</td>
          <td class="paramname"><em>detector_data_array</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">struct <a class="el" href="structsource__struct.html">source_struct</a> *&#160;</td>
          <td class="paramname"><em>source_data_array</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype"><a class="el" href="structulonglong2.html">ulonglong2</a> *&#160;</td>
          <td class="paramname"><em>materials_dose</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div>
<div class="memdoc">

<p>Initialize the image array, ie, set all pixels to zero Essentially, this function has the same effect as the command: "cutilSafeCall(cudaMemcpy(image_device, image, image_bytes, cudaMemcpyHostToDevice))";. </p>
<p>CUDA performs some initialization work the first time a GPU kernel is called. Therefore, calling a short kernel before the real particle tracking is performed may improve the accuracy of the timing measurements in the relevant kernel.</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">image</td><td>Pointer to the image array. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">pixels_per_image</td><td>Number of pixels in the image (ie, elements in the array). Main function to simulate x-ray tracks inside a voxelized geometry. Secondary electrons are not simulated (in photoelectric and Compton events the energy is locally deposited).</td></tr>
  </table>
  </dd>
</dl>
<p>The following global variables, in the GPU __constant__ memory are used: voxel_data_CONST, source_energy_data_CONST, detector_data_CONST, mfp_table_data_CONST.</p>
<dl class="params"><dt><b>Parameters:</b></dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">history_batch</td><td>Particle batch number (only used in the CPU version when CUDA is disabled!, the GPU uses the built-in variable threadIdx) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">num_p</td><td>Projection number in the CT simulation. This variable defines a specific angle and the corresponding source and detector will be used. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">histories_per_thread</td><td>Number of histories to simulate for each call to this function (ie, for GPU thread). </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">seed_input</td><td>Random number generator seed (the same seed is used to initialize the two MLCGs of RANECU). </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">voxel_mat_dens</td><td>Pointer to the voxel densities and material vector (the voxelized geometry), stored in GPU glbal memory. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">mfp_Woodcock_table</td><td>Two parameter table for the linear interpolation of the Woodcock mean free path (MFP) (stored in GPU global memory). </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">mfp_table_a</td><td>First element for the linear interpolation of the interaction mean free paths (stored in GPU global memory). </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">mfp_table_b</td><td>Second element for the linear interpolation of the interaction mean free paths (stored in GPU global memory). </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">rayleigh_table</td><td>Pointer to the table with the data required by the Rayleigh interaction sampling, stored in GPU global memory. </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">compton_table</td><td>Pointer to the table with the data required by the Compton interaction sampling, stored in GPU global memory. </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">image</td><td>Pointer to the image vector in the GPU global memory. </td></tr>
    <tr><td class="paramdir">[in,out]</td><td class="paramname">dose</td><td>Pointer to the array containing the 3D voxel dose (and its uncertainty) in the GPU global memory. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00135">135</a> of file <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html">MC-GPU_kernel_v1.3.cu</a>.</p>

<p>References <a class="el" href="MC-GPU__v1_83_8h_source.html#l00311">dose_ROI_x_max_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00229">linear_interp::e0</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01287">GCOa()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01181">GRAa()</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00229">linear_interp::ide</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00841">init_PRNG()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01033">locate_voxel()</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00062">MAX_MATERIALS</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00335">mfp_table_data_CONST</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00254">rayleigh_struct::pmax</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00965">ranecu()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00995">ranecu_double()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01103">rotate_double()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00626">source()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00482">tally_image()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l01547">tally_materials_dose()</a>, <a class="el" href="MC-GPU__kernel__v1_83_8cu_source.html#l00418">tally_voxel_energy_deposition()</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00132">float2::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::x</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00132">float2::y</a>, <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::y</a>, and <a class="el" href="MC-GPU__v1_83_8h_source.html#l00133">float3::z</a>.</p>

<p>Referenced by <a class="el" href="MC-GPU__v1_83_8cu_source.html#l00377">main()</a>.</p>

</div>
</div>
</div><!-- contents -->
</div>
  <div id="nav-path" class="navpath">
    <ul>
      <li class="navelem"><a class="el" href="MC-GPU__kernel__v1_83_8cu.html">MC-GPU_kernel_v1.3.cu</a>      </li>

    <li class="footer">Generated on Mon Dec 17 2012 13:26:39 for MC-GPU by
    <a href="http://www.doxygen.org/index.html">
    <img class="footer" src="doxygen.png" alt="doxygen"/></a> 1.7.6.1 </li>
   </ul>
 </div>


</body>
</html>
